[CUB] Fix DeviceAdjacentDifference for cuda::device_buffer iterators by Noperi0r · Pull Request #9861 · NVIDIA/cccl

Noperi0r · 2026-07-14T17:20:57Z

Description

cub::DeviceAdjacentDifference::SubtractLeftCopy fails to compile with cuda::device_buffer<T>::iterator because AgentDifference selects LoadIt through try_make_cache_modified_iterator_t but constructs it directly with LoadIt(input_it).

This change constructs load_it through try_make_cache_modified_iterator, matching the existing type-selection path. It also adds a regression test that exercises both the temporary-storage query and execution calls with cuda::device_buffer<int>::iterator.

Validation:

Focused regression test: 1/1 passed
Applicable AdjacentDifference lid_0 suite: 6/6 passed
Standalone ClangCUDA 21.1.8 and NVCC 12.9.86 reproducer/control matrix: 6/6 compile-and-run combinations passed
compute-sanitizer --tool memcheck: 0 errors
Targeted pre-commit hooks: passed
Raw-pointer AdjacentDifference SASS: byte-identical before and after

Checklist

New or existing tests cover these changes.
The documentation is up to date with these changes.

copy-pr-bot · 2026-07-14T17:21:00Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

coderabbitai · 2026-07-14T17:24:37Z

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

@coderabbitai resume to resume automatic reviews.
@coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

▶️ Resume reviews
🔍 Trigger review

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 5bfcb309-9743-4f0f-b93e-19d6a91490e7

📥 Commits

Reviewing files that changed from the base of the PR and between 8c1fa4f and 4ced5e9.

📒 Files selected for processing (1)

cub/test/catch2_test_device_adjacent_difference_substract_left.cu

🚧 Files skipped from review as they are similar to previous changes (1)

cub/test/catch2_test_device_adjacent_difference_substract_left.cu

📝 Walkthrough

Summary by CodeRabbit

Bug Fixes
- Fixed adjacent-difference correctness when using cached input iterators by adjusting how the cached-modified input iterator is created, ensuring proper data loading for subsequent adjacent-difference operations.
Tests
- Added a missing header needed to use CUDA device buffers in tests.
- Added a new device-buffer-based test for subtract-left copy, exercising execution on a CUDA stream and validating the expected output values.

Walkthrough

AgentDifference now creates cache-modified load iterators through the supported helper. A Catch2 test covers cuda::device_buffer input for SubtractLeftCopy and validates the subtraction output.

Changes

Adjacent Difference Iterator Compatibility

Layer / File(s)	Summary
Cache-modified iterator construction `cub/cub/agent/agent_adjacent_difference.cuh`	`AgentDifference` creates `load_it` through `try_make_cache_modified_iterator<Policy::LOAD_MODIFIER>`.
Device-buffer API regression coverage `cub/test/catch2_test_device_adjacent_difference_substract_left.cu`	The test includes `cuda/buffer` and validates `SubtractLeftCopy` with a `cuda::device_buffer` input and expected subtraction results.

Assessment against linked issues

Objective	Addressed	Explanation
Accept `cuda::device_buffer<T>::iterator` in `DeviceAdjacentDifference::SubtractLeftCopy`, including temporary-storage query and execution calls [`#9859`]	✅

Possibly related PRs

NVIDIA/cccl#9894: Changes contiguity detection in the same cache-modified iterator helper.

Suggested reviewers: pauleonix

_{Comment @coderabbitai help to get the list of available commands.}

bernhardmgruber · 2026-07-14T22:08:27Z

Important: Please move the new test into the file catch2_test_device_adjacent_difference_substract_left.cu. The _api.cu files cover tested API examples for the documentation.

Moved the regression test to
catch2_test_device_adjacent_difference_substract_left.cu
and restored catch2_test_device_adjacent_difference_api.cu to its previous scope.

The relocated focused test, target build, and targeted pre-commit checks pass locally.

bernhardmgruber · 2026-07-14T22:08:40Z

+
+  cuda::stream stream{cuda::devices[0]};
+  auto input = cuda::make_device_buffer<type>(stream, cuda::devices[0], {2, 5, 9, 14, 20});
+  c2h::device_vector<type> output(input.size());


Suggested change

c2h::device_vector<type> output(input.size());

c2h::device_vector<type> output(input.size(), thrust::no_init);

Applied thrust::no_init to the output allocation in the relocated test.

bernhardmgruber · 2026-07-14T22:09:05Z

+    == cub::DeviceAdjacentDifference::SubtractLeftCopy(
+      temporary_storage, temporary_bytes, input.begin(), output_it, input.size(), cuda::std::minus<>{}, stream.get()));
+
+  c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes);


Suggested change

c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes);

c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes, thrust::no_init);

Applied thrust::no_init to the temporary-storage allocation in the relocated test.

coderabbitai

Actionable comments posted: 1

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: de06a381-b09a-4138-9843-f913ccbe3dac

📥 Commits

Reviewing files that changed from the base of the PR and between 1b058db and facc0d5.

📒 Files selected for processing (1)

cub/test/catch2_test_device_adjacent_difference_substract_left.cu

Noperi0r · 2026-07-15T15:55:52Z

Addressed the review feedback in the latest commits:

Moved the cuda::device_buffer regression test from
catch2_test_device_adjacent_difference_api.cu to
catch2_test_device_adjacent_difference_substract_left.cu.
Restored the API test file to its previous scope.
Applied thrust::no_init to both the output and temporary-storage allocations.
Made the output_it raw-pointer binding const following the CodeRabbit coding-guideline suggestion. This only prevents reassignment of the local pointer variable and does not make the output buffer const.

The target build, focused regression test, and targeted pre-commit checks pass locally.

bernhardmgruber · 2026-07-15T16:02:02Z

+  void* temporary_storage     = nullptr;
+  std::size_t temporary_bytes = 0;
+  REQUIRE(
+    cudaSuccess
+    == cub::DeviceAdjacentDifference::SubtractLeftCopy(
+      temporary_storage, temporary_bytes, input.begin(), output_it, input.size(), cuda::std::minus<>{}, stream.get()));
+
+  c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes, thrust::no_init);
+  temporary_storage = thrust::raw_pointer_cast(temporary_buffer.data());
+  REQUIRE(
+    cudaSuccess
+    == cub::DeviceAdjacentDifference::SubtractLeftCopy(
+      temporary_storage, temporary_bytes, input.begin(), output_it, input.size(), cuda::std::minus<>{}, stream.get()));


Important: Just use the launch wrapper adjacent_difference_subtract_left instead of handling the temporary storage allocation yourself.

Updated the regression test to use adjacent_difference_subtract_left_copy, since this test covers SubtractLeftCopy. The launch wrapper now handles the temporary storage internally.

The focused target build, regression test, and targeted pre-commit checks pass locally.

bernhardmgruber · 2026-07-15T17:13:28Z

+
+  const c2h::host_vector<type> expected{2, 3, 4, 5, 6};
+  REQUIRE(output == expected);


Important: I think we do not to sync when we use a custom stream:

Suggested change

const c2h::host_vector<type> expected{2, 3, 4, 5, 6};

REQUIRE(output == expected);

stream.sync();

const c2h::host_vector<type> expected{2, 3, 4, 5, 6};

REQUIRE(output == expected);

Added stream.sync() immediately after the custom-stream launch and before the host-side output comparison.

The updated test builds and passes locally, along with git diff --check and the targeted pre-commit checks.

Fix DeviceAdjacentDifference for device_buffer iterators [skip-docs]

1b058db

Noperi0r requested a review from a team as a code owner July 14, 2026 17:20

Noperi0r requested a review from pauleonix July 14, 2026 17:20

github-project-automation Bot added this to CCCL Jul 14, 2026

github-project-automation Bot moved this to Todo in CCCL Jul 14, 2026

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jul 14, 2026

bernhardmgruber reviewed Jul 14, 2026

View reviewed changes

Move DeviceAdjacentDifference device_buffer test [skip-docs]

facc0d5

coderabbitai Bot reviewed Jul 15, 2026

View reviewed changes

Comment thread cub/test/catch2_test_device_adjacent_difference_substract_left.cu Outdated

Make AdjacentDifference output pointer binding const [skip-docs]

7aed3a4

bernhardmgruber reviewed Jul 15, 2026

View reviewed changes

Use launch wrapper in device_buffer regression test [skip-docs]

8c1fa4f

bernhardmgruber reviewed Jul 15, 2026

View reviewed changes

coderabbitai Bot mentioned this pull request Jul 16, 2026

Use thrust::is_contiguous_iterator_v over cuda::std::contiguous_iterator #9894

Open

Synchronize custom stream in device_buffer regression test [skip-docs]

4ced5e9

	c2h::device_vector<type> output(input.size());
	c2h::device_vector<type> output(input.size(), thrust::no_init);

	c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes);
	c2h::device_vector<std::uint8_t> temporary_buffer(temporary_bytes, thrust::no_init);


		const c2h::host_vector<type> expected{2, 3, 4, 5, 6};
		REQUIRE(output == expected);

Uh oh!

Conversation

Noperi0r commented Jul 14, 2026

Description

Checklist

Uh oh!

copy-pr-bot Bot commented Jul 14, 2026

Uh oh!

coderabbitai Bot commented Jul 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Reviews paused

Summary by CodeRabbit

Walkthrough

Changes

Assessment against linked issues

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Noperi0r commented Jul 15, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

coderabbitai Bot commented Jul 14, 2026 •

edited

Loading